{
 "cells": [
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "# PyOpenCL: Experimenting in IPython"
   ]
  },
  {
   "cell_type": "code",
   "execution_count": null,
   "metadata": {
    "collapsed": false
   },
   "outputs": [],
   "source": [
    "from __future__ import division\n",
    "import numpy as np\n",
    "import pyopencl as cl\n",
    "import pyopencl.array"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "Load the PyOpenCL IPython extension:"
   ]
  },
  {
   "cell_type": "code",
   "execution_count": 2,
   "metadata": {
    "collapsed": false
   },
   "outputs": [
    {
     "name": "stderr",
     "output_type": "stream",
     "text": [
      "/usr/lib/python3/dist-packages/IPython/utils/traitlets.py:504: DeprecationWarning: inspect.getargspec() is deprecated, use inspect.signature() instead\n",
      "  argspec = inspect.getargspec(c)\n"
     ]
    }
   ],
   "source": [
    "%load_ext pyopencl.ipython_ext"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "Create an OpenCL context and a command queue:"
   ]
  },
  {
   "cell_type": "code",
   "execution_count": 3,
   "metadata": {
    "collapsed": false
   },
   "outputs": [],
   "source": [
    "ctx = cl.create_some_context()\n",
    "queue = cl.CommandQueue(ctx)"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "## Using the kernel 'magic'\n",
    "\n",
    "Define an OpenCL kernel using the `%%cl_kernel` magic:"
   ]
  },
  {
   "cell_type": "code",
   "execution_count": 4,
   "metadata": {
    "collapsed": false
   },
   "outputs": [],
   "source": [
    "%%cl_kernel\n",
    "\n",
    "__kernel void sum_vector(__global const float *a,\n",
    "__global const float *b, __global float *c)\n",
    "{\n",
    "  int gid = get_global_id(0);\n",
    "  c[gid] = a[gid] + b[gid];\n",
    "}"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "This looks for `cl_ctx` or `ctx` in the user namespace to find a PyOpenCL context.\n",
    "\n",
    "Kernel names are automatically injected into the user namespace, so we can just use `sum_vector` from Python below.\n",
    "\n",
    "Now create some data to work on:"
   ]
  },
  {
   "cell_type": "code",
   "execution_count": 5,
   "metadata": {
    "collapsed": false
   },
   "outputs": [],
   "source": [
    "n = 10000\n",
    "\n",
    "a = cl.array.empty(queue, n, dtype=np.float32)\n",
    "a.fill(15)\n",
    "\n",
    "b_host = np.random.randn(n).astype(np.float32)\n",
    "b = cl.array.to_device(queue, b_host)\n",
    "\n",
    "c = cl.array.empty_like(a)"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "Run the kernel:"
   ]
  },
  {
   "cell_type": "code",
   "execution_count": 6,
   "metadata": {
    "collapsed": false
   },
   "outputs": [
    {
     "data": {
      "text/plain": [
       "<pyopencl.cffi_cl.Event at 0x7f48c1df4ba8>"
      ]
     },
     "execution_count": 6,
     "metadata": {},
     "output_type": "execute_result"
    }
   ],
   "source": [
    "sum_vector(queue, (n,), None, a.data, b.data, c.data)"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "Check the result using `numpy` operations:"
   ]
  },
  {
   "cell_type": "code",
   "execution_count": 7,
   "metadata": {
    "collapsed": false
   },
   "outputs": [],
   "source": [
    "assert (c.get() == b_host + 15).all()"
   ]
  }
 ],
 "metadata": {
  "kernelspec": {
   "display_name": "Python 3",
   "language": "python",
   "name": "python3"
  },
  "language_info": {
   "codemirror_mode": {
    "name": "ipython",
    "version": 3
   },
   "file_extension": ".py",
   "mimetype": "text/x-python",
   "name": "python",
   "nbconvert_exporter": "python",
   "pygments_lexer": "ipython3",
   "version": "3.5.0+"
  }
 },
 "nbformat": 4,
 "nbformat_minor": 0
}
